#include <cuda_runtime.h>

// 定义一个简单的 kernel 函数
__global__ void simple_slice_async_example(float* input, float* output, int size) {
    // 定义 L1 缓冲区
    __valigned__ float l1_buffer[256];

    // 创建 mdspan 对象
    tops::mdspan l3_input(tops::Global, input, {size});
    tops::mdspan l1_output(tops::Private, l1_buffer, {256});

    // 创建 DTE 上下文
    tops_dte_ctx_t dte;
    dte.init();

    // 定义事件
    tops::event event;

    // 使用 slice_async 从 L3 加载数据到 L1
    event = tops::slice_async(dte, l1_output, l3_input, {0});

    // 等待数据加载完成
    tops::wait(event);

    // 在 L1 上进行简单计算（例如加 1）
    // for (int i = 0; i < 256; i++) {
    //     l1_buffer[i] += 1.0f;
    // }

    // 将结果写回 L3
    event = tops::deslice_async(dte, l3_input, l1_output, {0});
    tops::wait(event);

    // 销毁 DTE
    dte.destroy();
}

void run_example(float* dev_input, float* dev_output, int size) {
    // 定义 kernel 启动配置
    static const size_t blocks = 1;
    static const size_t threads = 1;

    // 启动 kernel
    simple_slice_async_example<<<blocks, threads>>>(dev_input, dev_output, size);
}
